/*
目的:
验证, cuda块线程中:
    每个cuda线程, 局部变量是独立的
*/

#include <stdio.h>
#include <time.h>
#include <math.h>
#include <cuda_runtime.h>

float gpu_clockRate = 0.0;
#define GPU_TASKS 2
// cuda核函数: 并行计算数组元素的开方 和
__global__ void parallel_cal_array_sqrt_sum(float* array, int size, float* result, float* record)
{
    // cuda线程id
    const int tid = threadIdx.x;

    // 要计算的元素个数
    const int cal_size = size / GPU_TASKS;
    float sum = 0.0;

    for (int i = tid * cal_size; i < (tid + 1) * cal_size; i++) {
        // cuda函数 --> sqrtf: float开方函数, sqrt: double开方函数
        sum += array[i];
        record[i] = sum;
    }
    result[tid] = sum;
}

float gpu_parallel_test(float* array, int size)
{
    // host 变量
    float result[GPU_TASKS];
    float record[6];
    clock_t clock_consum = 0;

    // device 变量
    float* d_array, * d_result, * d_record;
    clock_t* d_clock_consum;

    // cauda device申请内存
    cudaMalloc((void**)&d_array, size * sizeof(float));
    cudaMalloc((void**)&d_record, size * sizeof(float));
    cudaMalloc((void**)&d_result, GPU_TASKS * sizeof(float));
    cudaMalloc((void**)&d_clock_consum, sizeof(clock_t));

    // 将host数据拷贝到device
    cudaMemcpy(d_array, array, size * sizeof(float), cudaMemcpyHostToDevice);

    // 调用核函数
    parallel_cal_array_sqrt_sum << <1, GPU_TASKS, 0 >> > (d_array, size, d_result, d_record);

    // 将device数据拷贝到host
    cudaMemcpy(&result, d_result, GPU_TASKS * sizeof(float), cudaMemcpyDeviceToHost);
    cudaMemcpy(&clock_consum, d_clock_consum, sizeof(clock_t), cudaMemcpyDeviceToHost);
    cudaMemcpy(record, d_record, size * sizeof(float), cudaMemcpyDeviceToHost);

    // 释放device内存
    cudaFree(d_array);
    cudaFree(d_result);
    cudaFree(d_clock_consum);

    float sum = 0.0;
    for (int i = 0; i < GPU_TASKS; i++) {
        sum += result[i];
    }
    printf("gpu 并行计算结果:%f, 耗时:%f\n", sum, clock_consum / gpu_clockRate);

    // 打印 cuda 线程共享内存record
    for (int i = 0; i < size; i++) {
        printf("record[%d]:%f\n", i, record[i]);
    }

    return sum;
}

float cpu_test(float* array, int size)
{
    float result = 0.0;
    clock_t start = clock();
    for (int i = 0; i < size; i++) {
        result += array[i];
    }
    clock_t end = clock();

    printf("cpu 计算结果:%f, 耗时:%f\n", result, (end - start) / (float)CLOCKS_PER_SEC);
    return result;
}

void init_gpu_clockRate()
{
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    gpu_clockRate = prop.clockRate * 1000.0;
}

int main()
{
    init_gpu_clockRate();

    // 测试数据
    float array[] = { 1.0, 2.0, 3.0, 4.0, 5.0, 6.0 };
    int size = sizeof(array) / sizeof(float);

    gpu_parallel_test(array, size);
    cpu_test(array, size);

    return 0;
}

/* 运行结果:
gpu 并行计算结果:21.000000, 耗时:0.000000
record[0]:1.000000
record[1]:3.000000
record[2]:6.000000
record[3]:4.000000
record[4]:9.000000
record[5]:15.000000
cpu 计算结果:21.000000, 耗时:0.000001


测试结论:
每个cuda线程, 局部变量是独立的

 */